#ifndef _PDIST_KERNEL_H_
#define _PDIST_KERNEL_H_

// pdist - using 256 threads per block, partially coalesced
#define BLOCK_DIM 256

__global__ void pdist_gpu(float *output, float* input, int size_x, int size_y)
{
	extern __shared__ float temp[];
	float tmp = 0;
	float s = 0;
	unsigned int yIndex = blockIdx.y * BLOCK_DIM + threadIdx.y;

	// loop through the other elements
	for (int y = 0; y < size_y; y++)
	{
		s = 0;

		for (int i = 0; i <= size_x/BLOCK_DIM; i++)
		{
			// each thread loads one element of the current row
			if (threadIdx.y+i*BLOCK_DIM < size_x)
				temp[threadIdx.y] = input[y + (threadIdx.y+i*BLOCK_DIM)*size_y];
		}

		__syncthreads();

		// compute the distance
		for (int i = 0; i < size_x && yIndex < size_y; i++)
		{
			tmp = temp[i] - input[yIndex + i*size_y];
			s += tmp*tmp;
		}

		if (yIndex < size_y)
			output[yIndex + y*size_y] = sqrtf(s);

		__syncthreads();
	}

}

#endif // _PDIST_KERNEL_H_
